#include<stdio.h>
#include <cutil_inline.h>
#include<assert.h>
#include<stdlib.h>
#include<time.h>

#define ARRAY_SIZE 8388608
//#define ARRAY_SIZE 16777216
#define THREADS_PER_BLOCK 256

#define NBLOCKS ARRAY_SIZE / THREADS_PER_BLOCK 
#define ITER 10


__global__ void
execute(int* const dData)
{
    const int thread = blockIdx.x *  blockDim.x + threadIdx.x;
    dData[thread] += 10;
}


__global__ void
dummy()
{
    __shared__ int  temp;
    temp = 3;
    temp = temp + 1;
}

int
main()
{
    const unsigned size = sizeof(int) * ARRAY_SIZE;
    int* const data = (int*)malloc(size);
    

    int* dData = 0;
    cutilSafeCall(cudaMalloc((void**)&dData, size));

    dim3 threads(THREADS_PER_BLOCK);
    dim3 blocks(NBLOCKS);
    dummy <<< blocks, threads >>> ();
    

     for (unsigned i = 0; i < ARRAY_SIZE; ++i) {
          data[i] = 0;
     }

    
     cudaEvent_t startEvent, stopEvent;
     cutilSafeCall( cudaEventCreate(&startEvent));
     cutilSafeCall( cudaEventCreate(&stopEvent) );


     cudaEventRecord(startEvent, 0);
	 for (unsigned k = 0; k < ITER; ++k) {
         cutilSafeCall(cudaThreadSynchronize());

		 cutilSafeCall(cudaMemcpy(dData, data, size,
				 cudaMemcpyHostToDevice));
		 execute <<< blocks, threads >>> (dData);
		 cutilSafeCall(cudaMemcpy(data, dData, size,
				 cudaMemcpyDeviceToHost));
	 }

     cudaEventRecord(stopEvent, 0);
     cudaEventSynchronize(stopEvent);
     float elapsedTime = 0;
     cutilSafeCall( cudaEventElapsedTime(&elapsedTime, 
                    startEvent, stopEvent));
     cudaEventDestroy(startEvent);
     cudaEventDestroy(stopEvent);


     printf("Processing time for synch %f\n", elapsedTime); 
     cutilSafeCall(cudaFree(dData));
      

     for (unsigned i = 0; i < ARRAY_SIZE; ++i) {
         assert(data[i] == ITER * 10);
     }

     free(data);
}

